// CUDA Device Query
 
#include <stdio.h>
#include <thrust/device_vector.h>
#include <iostream>
#include <vector>
#include <pthread.h>
#include <sys/time.h> 
using namespace std;
 
// Print device properties
void printDevProp(cudaDeviceProp devProp)
{
    printf("Major revision number:         %d\n",  devProp.major);
    printf("Minor revision number:         %d\n",  devProp.minor);
    printf("Name:                          %s\n",  devProp.name);
    printf("Total global memory:           %u\n",  devProp.totalGlobalMem);
    printf("Total shared memory per block: %u\n",  devProp.sharedMemPerBlock);
    printf("Total registers per block:     %d\n",  devProp.regsPerBlock);
    printf("Warp size:                     %d\n",  devProp.warpSize);
    printf("Maximum memory pitch:          %u\n",  devProp.memPitch);
    printf("Maximum threads per block:     %d\n",  devProp.maxThreadsPerBlock);
    for (int i = 0; i < 3; ++i)
    printf("Maximum dimension %d of block:  %d\n", i, devProp.maxThreadsDim[i]);
    for (int i = 0; i < 3; ++i)
    printf("Maximum dimension %d of grid:   %d\n", i, devProp.maxGridSize[i]);
    printf("Clock rate:                    %d\n",  devProp.clockRate);
    printf("Total constant memory:         %u\n",  devProp.totalConstMem);
    printf("Texture alignment:             %u\n",  devProp.textureAlignment);
    printf("Concurrent copy and execution: %s\n",  (devProp.deviceOverlap ? "Yes" : "No"));
    printf("Number of multiprocessors:     %d\n",  devProp.multiProcessorCount);
    printf("Kernel execution timeout:      %s\n",  (devProp.kernelExecTimeoutEnabled ? "Yes" : "No"));
    return;
}
 
struct Test1{
	__device__
	void operator()(int& b){
		double k = 10;
		for(int i = 0; i < 2; i++) k += sin(k);
		b = 99;
	}
}; 
 
__global__ void Test2(int* array){
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	double k = 10;
	for(int i = 0; i < 2; i++) k += sin(k);
	array[idx] = 99;
} 
 
struct Obj{
	vector<int> a;
	int k;
	bool temp_func(int a){
		return a > 0;
	}
}; 
 
//thrust::device_vector<Obj> mix_vec;
 
int size = 1000000; 
 
thrust::device_vector<int>* a;
int* b;
 
struct thread_id{
	int id;
	cudaStream_t stream;
	cudaEvent_t start;
	cudaEvent_t stop;
	
	pthread_cond_t work;
	pthread_mutex_t mutex;
	
	pthread_cond_t work_done;
	pthread_mutex_t done_mutex;
	
	bool ready1;
	bool ready2;
	
	int shared_temp;
	bool terminate;
}; 



void *copy(void* id){
	thread_id& task = *((thread_id*)id);
	cudaEventRecord(task.start, task.stream);
	// while(true){
		// pthread_mutex_lock(&(task.mutex));
		// while(!task.ready1)
			// pthread_cond_wait(&(task.work),&(task.mutex));
		// task.ready1 = false;
		// pthread_mutex_unlock(&(task.mutex));
		// if(task.terminate) break;
	
		cout << "thread " << task.id << " is running" << endl;
		for(int i = 0; i < 2000; i++){
	
			if(task.id == 0){
				// *a = *b;
				cudaMemcpy(thrust::raw_pointer_cast(a->data()),b,
					size*sizeof(int),cudaMemcpyHostToDevice);
				// cudaMemcpyAsync(thrust::raw_pointer_cast(a->data()),b,size*sizeof(int),cudaMemcpyHostToDevice,task.stream);
			}
			else{
				/*cudaMemcpy(thrust::raw_pointer_cast(a->data()),b,
					size*sizeof(int),cudaMemcpyHostToDevice);*/
				Test2<<<size/1000,1000,0,task.stream>>>(thrust::raw_pointer_cast(a->data()));
				// thrust::for_each(a->begin(),a->end(),Test1());
			}
		}
		
		// pthread_mutex_lock(&(task.done_mutex));
		// pthread_cond_signal(&(task.work_done));
		// task.ready2 = true;
		// pthread_mutex_unlock(&(task.done_mutex));
	// }
	cudaEventRecord(task.stop, task.stream);
	return NULL;
} 


 
int main()
{
	a = new thrust::device_vector<int>(size,0);
	cudaMallocHost((void**)&b,size*sizeof(int));
	
	cudaStream_t str[2];
	for(int i = 0; i < 2; i++) cudaStreamCreate(&(str[i]));
	
	struct timeval start, finish;
	gettimeofday(&start,NULL);
	
	// #pragma omp parallel for
	for(int i = 0; i < 2; i++){
		if(i == 0){
		int idx = 0;
		while(idx < 2000){
			cudaMemcpyAsync(thrust::raw_pointer_cast(a->data()),b,
				size*sizeof(int),cudaMemcpyHostToDevice,str[1]);
			idx++;
		}
		}else{
		
		int idx = 0;
		while(idx < 2000){
			Test2<<<size/1000,1000,0,str[1]>>>(thrust::raw_pointer_cast(a->data()));
			idx++;
		}
		
		}
	}
	
	
	// for(int j = 0; j < 2000; j++){
		// cudaMemcpyAsync(thrust::raw_pointer_cast(a->data()),b,
			// size*sizeof(int),cudaMemcpyHostToDevice,str[1]);
		// Test2<<<size/1000,1000,0,str[0]>>>(thrust::raw_pointer_cast(a->data()));
	// }
	
	cudaDeviceSynchronize();
	
	gettimeofday(&finish,NULL);
	double t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	double t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "cpu takes time: " << t2 - t1 << endl;
	exit(1);
	
    // Number of CUDA devices
    // int devCount;
    // cudaGetDeviceCount(&devCount);
    // printf("CUDA Device Query...\n");
    // printf("There are %d CUDA devices.\n", devCount);
 
    // Iterate through devices
    // for (int i = 0; i < devCount; ++i)
    // {
    //    Get device properties
        // printf("\nCUDA Device #%d\n", i);
        // cudaDeviceProp devProp;
        // cudaGetDeviceProperties(&devProp, i);
        // printDevProp(devProp);
    // }
 
    // printf("\nPress any key to exit...");
    // char c;
    // scanf("%c", &c);
	
	// b.resize(size,2);
	
	thread_id id_vec[2];
	for(int i = 0; i < 2; i++){
		id_vec[i].id = i;
		cudaEventCreate(&(id_vec[i].start));
		cudaEventCreate(&(id_vec[i].stop));
		cudaStreamCreate(&(id_vec[i].stream));
		pthread_mutex_init(&(id_vec[i].mutex),NULL);
		pthread_cond_init(&(id_vec[i].work),NULL);
		pthread_mutex_init(&(id_vec[i].done_mutex),NULL);
		pthread_cond_init(&(id_vec[i].work_done),NULL);
		id_vec[i].ready1 = false;
		id_vec[i].ready2 = false;
		id_vec[i].terminate = false;
	}
	
	gettimeofday(&start,NULL);
	
	pthread_t threads[2];
	for(int i = 0; i < 2; i++){
		pthread_create(&threads[i],NULL,copy,(void*)&(id_vec[i]));
	}
	
	// cout << "reach here" << endl;
	// for(int i = 0; i < 2; i++){
		// pthread_mutex_lock(&(id_vec[i].mutex));
		// id_vec[i].ready1 = true;
		// pthread_cond_signal(&(id_vec[i].work));
		// pthread_mutex_unlock(&(id_vec[i].mutex));
	// }
	
	// cout << "reach xxxx" << endl;
	
	// for(int i = 0; i < 2; i++){
		// pthread_mutex_lock(&(id_vec[i].done_mutex));
		// while(!id_vec[i].ready2)
			// pthread_cond_wait(&(id_vec[i].work_done),&(id_vec[i].done_mutex));
		// pthread_mutex_unlock(&(id_vec[i].done_mutex));
	// }
	
	// for(int i = 0; i < 2; i++){
		// pthread_mutex_lock(&(id_vec[i].mutex));
		// id_vec[i].ready1 = true;
		// id_vec[i].terminate = true;
		// pthread_cond_signal(&(id_vec[i].work));
		// pthread_mutex_unlock(&(id_vec[i].mutex));
	// }
	// while(id_vec[1].work != 0 && id_vec[0].work != 0){
		// sleep(5);
	// }
	// id_vec[0].terminate = 1;
	// id_vec[1].terminate = 1;
	
	for(int i = 0; i < 2; i++){
		pthread_join(threads[i],NULL);
	}
	cudaDeviceSynchronize();
	
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "cpu takes time: " << t2 - t1 << endl;
	
	for(int i = 0; i < 2; i++){
		float time;
		cudaEventElapsedTime(&time, id_vec[i].start, id_vec[i].stop);
		cout << "Device " << i << " takes time: " << time << endl;
	}
	
	for(int i = 0; i < 100; i++) cout << (*a)[i] << " ";
	cout << endl;
	exit(1);
 
	// int* test[2];
	// cudaMalloc((void**)&test[0],size*sizeof(int));
	// cudaMalloc((void**)&test[1],size*sizeof(int));
	thrust::device_vector<int> *test[2];
	test[0] = new thrust::device_vector<int>(size,0);
	test[1] = new thrust::device_vector<int>(size,0);
	
	int* h_test;
	cudaMallocHost((void **)&h_test, size*sizeof(int));
	for(int i = 0; i < size; i++) h_test[i] = 5;
	// vector<int> h_test(size,5);
	
	
	cout << "finished" << endl;
	cudaStream_t stream[2];
	cudaStreamCreate(&stream[0]);cudaStreamCreate(&stream[1]);
	
	// cudaMemcpyAsync(test[0],h_test,size*sizeof(int),cudaMemcpyHostToDevice,stream[0]);
	// cudaMemcpyAsync(test[1],h_test,size*sizeof(int),cudaMemcpyHostToDevice,stream[0]);
	// Test2<<<size/1000,1000,0,stream[1]>>>(test[1]);
	
	cudaMemcpyAsync(thrust::raw_pointer_cast(test[0]->data()),h_test,size*sizeof(int),cudaMemcpyHostToDevice,stream[0]);
	cudaMemcpyAsync(thrust::raw_pointer_cast(test[1]->data()),h_test,size*sizeof(int),cudaMemcpyHostToDevice,stream[0]);
	
	Test2<<<1,32,0,stream[1]>>>(thrust::raw_pointer_cast(test[1]->data()));
	//thrust::for_each(test[1]->begin(),test[1]->begin()+32,Test1());
	cudaDeviceSynchronize();
	
	// cudaMemcpy(h_test,test[1],size*sizeof(int),cudaMemcpyDeviceToHost);
	
	for(int i = 0; i < 15; i++) cout << (*test[1])[i] << " ";
	cout << endl;
 
    return 0;
}